{
  "cells": [
    {
      "attachments": {},
      "cell_type": "markdown",
      "metadata": {},
      "source": [
        "(vta-get-started)=\n",
        "# VTA 入门\n",
        "\n",
        "**原作者**: [Thierry Moreau](https://homes.cs.washington.edu/~moreau/)|修改 [xinetzone](https://github.com/xinetzone)\n",
        "\n",
        "这是关于如何使用 TVM 编程 VTA 设计的入门教程。\n",
        "\n",
        "在本教程中，将演示在 VTA 设计的向量 ALU 上实现向量加法的基本 TVM 工作流。此过程包括将计算 lower 到低级加速器运算所需的特定调度变换。\n",
        "\n",
        "首先，需要导入 TVM，这是深度学习优化编译器。还需要导入 VTA python 包，其中包含针对 TVM 的 VTA 特定扩展，以实现 VTA 设计。"
      ]
    },
    {
      "cell_type": "code",
      "execution_count": 1,
      "metadata": {
        "tags": [
          "remove-cell"
        ]
      },
      "outputs": [],
      "source": [
        "import set_env"
      ]
    },
    {
      "cell_type": "code",
      "execution_count": 2,
      "metadata": {
        "collapsed": false
      },
      "outputs": [],
      "source": [
        "import os\n",
        "import tvm\n",
        "from tvm import te\n",
        "import vta\n",
        "from tvm.ir.module import IRModule\n",
        "import numpy as np"
      ]
    },
    {
      "cell_type": "markdown",
      "metadata": {},
      "source": [
        "## 加载 VTA 参数\n",
        "\n",
        "VTA 是模块化和可定制的设计。因此，用户可以自由地修改影响硬件设计布局的高级硬件参数。这些参数在 [tvm/vta/vta_hw/config/vta_config.json](https://github.com/xinetzone/tvm/tree/dev/vta/vta_hw/config/vta_config.json) 中通过它们的 `log2` 值指定。\n",
        "这些 VTA 参数可以通过 `vta.get_env` 函数加载。\n",
        "\n",
        "最后，TVM 目标也在 `vta_config.json` 文件中指定。当设置为 `sim` 时，执行将发生在 VTA 仿真器行为内。如果您想在 Pynq FPGA 开发平台上运行本教程，请遵循 *VTA 基于 Pynq 的测试设置* 指南。"
      ]
    },
    {
      "cell_type": "code",
      "execution_count": 3,
      "metadata": {
        "collapsed": false
      },
      "outputs": [],
      "source": [
        "env = vta.get_env()"
      ]
    },
    {
      "cell_type": "markdown",
      "metadata": {},
      "source": [
        "## FPGA 编程\n",
        "\n",
        "当针对 Pynq FPGA 开发板时，需要使用 VTA bitstream 配置该板。\n",
        "\n",
        "需要 TVM RPC 模块和 VTA 仿真器模块："
      ]
    },
    {
      "cell_type": "code",
      "execution_count": 4,
      "metadata": {},
      "outputs": [],
      "source": [
        "from vta.testing import simulator # 此处一定要有"
      ]
    },
    {
      "cell_type": "markdown",
      "metadata": {},
      "source": [
        "```{warning}\n",
        "若 vta 是 `sim` 模式，一定要载入 `simulator` 模块，否则会触发异常。\n",
        "```\n",
        "\n",
        "从操作系统环境中读取 Pynq RPC 主机 IP 地址和端口号："
      ]
    },
    {
      "cell_type": "code",
      "execution_count": 5,
      "metadata": {},
      "outputs": [],
      "source": [
        "host = os.environ.get(\"VTA_RPC_HOST\", \"192.168.2.99\")\n",
        "port = int(os.environ.get(\"VTA_RPC_PORT\", \"9091\"))"
      ]
    },
    {
      "cell_type": "markdown",
      "metadata": {},
      "source": [
        "在 Pynq 上配置 bitstream 和运行时系统，以匹配 `vta_config.json` 文件指定的 VTA 配置。"
      ]
    },
    {
      "cell_type": "code",
      "execution_count": 6,
      "metadata": {},
      "outputs": [
        {
          "data": {
            "text/plain": [
              "'sim'"
            ]
          },
          "execution_count": 6,
          "metadata": {},
          "output_type": "execute_result"
        }
      ],
      "source": [
        "env.TARGET"
      ]
    },
    {
      "cell_type": "code",
      "execution_count": 7,
      "metadata": {
        "collapsed": false
      },
      "outputs": [],
      "source": [
        "if env.TARGET in [\"pynq\", \"de10nano\"]:\n",
        "    # 确保使用 RPC=1 编译 TVM\n",
        "    assert tvm.runtime.enabled(\"rpc\")\n",
        "    remote = tvm.rpc.connect(host, port)\n",
        "\n",
        "    # 重新配置 JIT runtime\n",
        "    vta.reconfig_runtime(remote)\n",
        "\n",
        "    # 用预编译的 VTA bitstream 编程 FPGA。\n",
        "    # 通过将 path 传递给 bitstream 文件而不是 None，\n",
        "    # 您可以使用自定义 bitstream 编程 FPGA。\n",
        "    vta.program_fpga(remote, bitstream=None)\n",
        "\n",
        "# 在仿真模式中，在本地托管 RPC 服务器。\n",
        "elif env.TARGET in (\"sim\", \"tsim\", \"intelfocl\"):\n",
        "    remote = tvm.rpc.LocalSession()\n",
        "\n",
        "    if env.TARGET in [\"intelfocl\"]:\n",
        "        # program intelfocl aocx\n",
        "        vta.program_fpga(remote, bitstream=\"vta.bitstream\")"
      ]
    },
    {
      "attachments": {},
      "cell_type": "markdown",
      "metadata": {},
      "source": [
        "## 准备数据\n",
        "\n",
        "为了验证计算的正确性，需要准备一些数据, 随机初始化 A 和 B 数组，`int` 范围为 $(-128, 128]$："
      ]
    },
    {
      "cell_type": "code",
      "execution_count": 8,
      "metadata": {},
      "outputs": [],
      "source": [
        "m = 16\n",
        "n = 1024\n",
        "A_orig = np.random.randint(-128, 128, size=(m, n)).astype(env.acc_dtype)\n",
        "B_orig = np.random.randint(-128, 128, size=(m, n)).astype(env.acc_dtype)"
      ]
    },
    {
      "attachments": {},
      "cell_type": "markdown",
      "metadata": {},
      "source": [
        "为了适应于 VTA 设备，需要应用 packing，将 A 和 B 数组从 2D 到 4D packed layout："
      ]
    },
    {
      "cell_type": "code",
      "execution_count": 9,
      "metadata": {},
      "outputs": [],
      "source": [
        "A_packed = A_orig.reshape(m//env.BATCH, env.BATCH, n//env.BLOCK_OUT, env.BLOCK_OUT).transpose((0, 2, 1, 3))\n",
        "B_packed = B_orig.reshape(m//env.BATCH, env.BATCH, n//env.BLOCK_OUT, env.BLOCK_OUT).transpose((0, 2, 1, 3))"
      ]
    },
    {
      "cell_type": "markdown",
      "metadata": {},
      "source": [
        "## 计算声明\n",
        "\n",
        "第一步，需要描述计算。TVM 采用张量语义，每个中间结果表示为多维数组。用户需要描述生成输出张量的计算规则。\n",
        "\n",
        "在这个例子中，描述了向量加法，它需要多个计算阶段，如下面的数据流程图所示。\n",
        "\n",
        "- 首先，描述存在于 main memory 中的输入张量 `A` 和 `B`。\n",
        "- 其次，需要声明中间张量 `A_buf` 和 `B_buf`，它们将位于 VTA 的 on-chip buffers 中。有了这个额外的计算阶段，就可以显式地分阶段进行 cached 的读写操作。\n",
        "- 第三，描述向量加法运算：`C_buf = A_buf + B_buf`。\n",
        "- 最后的运算是强制转换并复制回 DRAM，到结果张量 `C` 中。\n",
        "\n",
        "\n",
        "```{image} images/vadd_dataflow.png\n",
        ":align: center\n",
        "```"
      ]
    },
    {
      "attachments": {},
      "cell_type": "markdown",
      "metadata": {},
      "source": [
        "### Input 占位符\n",
        "\n",
        "以平铺（tiled）数据格式描述占位符张量 `A` 和 `B`，以匹配 VTA 向量 ALU 施加的数据布局要求。\n",
        "\n",
        "对于 VTA 的一般用途的运算，如 ALU 加法，tile 大小为 `(env.BATCH, env.BLOCK_OUT)`。维度在 `vta_config.json` 配置文件中指定，默认设置为 (1, 16) 向量。"
      ]
    },
    {
      "cell_type": "code",
      "execution_count": 10,
      "metadata": {},
      "outputs": [
        {
          "data": {
            "text/plain": [
              "(16, 64, 1, 16)"
            ]
          },
          "execution_count": 10,
          "metadata": {},
          "output_type": "execute_result"
        }
      ],
      "source": [
        "# 输出通道因子 m -总共 64 x 16 = 1024 输出通道\n",
        "_m = n//env.BLOCK_OUT\n",
        "# Batch 因子 o - 总共 16 x 1 = 1\n",
        "_o = m//env.BATCH\n",
        "\n",
        "# VTA 向量数据 shape\n",
        "shape = (_o, _m, env.BATCH, env.BLOCK_OUT)\n",
        "shape"
      ]
    },
    {
      "cell_type": "markdown",
      "metadata": {},
      "source": [
        "查看 {data}`env.acc_dtype` 和 {data}`env.inp_dtype`："
      ]
    },
    {
      "cell_type": "code",
      "execution_count": 11,
      "metadata": {},
      "outputs": [
        {
          "data": {
            "text/plain": [
              "('int32', 'int8', 'int8', 'int8')"
            ]
          },
          "execution_count": 11,
          "metadata": {},
          "output_type": "execute_result"
        }
      ],
      "source": [
        "env.acc_dtype, env.inp_dtype, env.out_dtype, env.wgt_dtype"
      ]
    },
    {
      "cell_type": "markdown",
      "metadata": {},
      "source": [
        "此外，A 和 B 的数据类型也需要匹配 `env.acc_dtype`，由 `vta_config.json` 文件设置为 32 位整型。"
      ]
    },
    {
      "cell_type": "code",
      "execution_count": 12,
      "metadata": {
        "collapsed": false
      },
      "outputs": [
        {
          "data": {
            "text/plain": [
              "Tensor(shape=[16, 64, 1, 16], op.name=A)"
            ]
          },
          "execution_count": 12,
          "metadata": {},
          "output_type": "execute_result"
        }
      ],
      "source": [
        "# 平铺 A, B 占位符张量数据\n",
        "A = te.placeholder(shape, name=\"A\", dtype=env.acc_dtype)\n",
        "B = te.placeholder(shape, name=\"B\", dtype=env.acc_dtype)\n",
        "A"
      ]
    },
    {
      "cell_type": "markdown",
      "metadata": {},
      "source": [
        "### Copy Buffers\n",
        "\n",
        "硬件加速器的特点之一是，必须对 on-chip memory 进行显式管理。这意味着需要描述中间张量 `A_buf` 和 `B_buf`，它们可以具有与原始占位符张量 `A` 和 `B` 不同的内存作用域。\n",
        "\n",
        "稍后在调度阶段，可以告诉编译器 `A_buf` 和 `B_buf` 将存在于 VTA 的 on-chip buffer（SRAM）中，而 `A` 和 `B` 将存在于 main memory（DRAM）中。将 A_buf 和 B_buf 描述为恒等函数计算的运算结果。这可以被编译器解释为 cached 的读运算。"
      ]
    },
    {
      "cell_type": "code",
      "execution_count": 13,
      "metadata": {
        "collapsed": false
      },
      "outputs": [
        {
          "data": {
            "text/plain": [
              "Tensor(shape=[16, 64, 1, 16], op.name=A_buf)"
            ]
          },
          "execution_count": 13,
          "metadata": {},
          "output_type": "execute_result"
        }
      ],
      "source": [
        "# A copy buffer\n",
        "A_buf = te.compute(shape, lambda *i: A[i], \"A_buf\")\n",
        "# B copy buffer\n",
        "B_buf = te.compute(shape, lambda *i: B[i], \"B_buf\")\n",
        "A_buf"
      ]
    },
    {
      "attachments": {},
      "cell_type": "markdown",
      "metadata": {},
      "source": [
        "### ALU 加法\n",
        "\n",
        "现在可以用另一个 `compute` 运算来描述向量加法结果张量 `C`。`compute` 函数采用张量的形状，以及描述张量每个位置的计算规则的 lambda 函数。\n",
        "\n",
        "此阶段没有计算发生，因为只是声明了计算应该如何完成。"
      ]
    },
    {
      "cell_type": "code",
      "execution_count": 14,
      "metadata": {
        "collapsed": false
      },
      "outputs": [],
      "source": [
        "# 描述 VTA 中的 ALU 加法\n",
        "fcompute = lambda *i: A_buf[i].astype(env.acc_dtype) + B_buf[i].astype(env.acc_dtype)\n",
        "C_buf = te.compute(shape, fcompute, name=\"C_buf\")"
      ]
    },
    {
      "cell_type": "markdown",
      "metadata": {},
      "source": [
        "### Casting 结果\n",
        "\n",
        "计算完成后，需要将 VTA 计算的结果发送回主存储器（main memory）\n",
        "\n",
        "```{admonition} 内存存储限制\n",
        ":class: alert alert-info\n",
        "\n",
        "VTA 的特点之一是，它只支持窄化（narrow） `env.inp_dtype` 数据类型格式的 DRAM 存储。这让我们能够减少内存传输的数据 footprint（详见基本矩阵乘法的例子）。\n",
        "```\n",
        "\n",
        "对窄化的输入激活数据格式执行最后一个 typecast 运算。"
      ]
    },
    {
      "cell_type": "code",
      "execution_count": 15,
      "metadata": {
        "collapsed": false
      },
      "outputs": [],
      "source": [
        "# 转换为输出类型，并发送到 main memory\n",
        "fcompute = lambda *i: C_buf[i].astype(env.inp_dtype)\n",
        "C = te.compute(shape, fcompute, name=\"C\")"
      ]
    },
    {
      "cell_type": "markdown",
      "metadata": {},
      "source": [
        "这就结束了本教程的计算声明部分。"
      ]
    },
    {
      "cell_type": "markdown",
      "metadata": {},
      "source": [
        "## 调度计算\n",
        "\n",
        "虽然上面的几行描述了计算规则，但我们可以通过许多方式得到 `C`。TVM 要求用户提供名为 **调度** （*schedule*） 的计算实现。\n",
        "\n",
        "调度是对原始计算的一组变换，它在不影响正确性的情况下变换计算的实现。这个简单的 VTA 编程教程旨在演示基本的调度变换，将原始调度映射到 VTA 硬件原语（primitives）。\n",
        "\n",
        "### 默认调度\n",
        "\n",
        "在构造了调度之后，默认情况下，调度会以如下方式计算 `C`："
      ]
    },
    {
      "cell_type": "code",
      "execution_count": 16,
      "metadata": {},
      "outputs": [
        {
          "data": {
            "text/html": [
              "<div class=\"highlight\" style=\"background: \"><pre style=\"line-height: 125%;\"><span></span><span style=\"color: #007979; font-style: italic\"># from tvm.script import ir as I</span>\n",
              "<span style=\"color: #007979; font-style: italic\"># from tvm.script import tir as T</span>\n",
              "\n",
              "<span style=\"color: #A2F\">@I</span><span style=\"color: #A2F; font-weight: bold\">.</span>ir_module\n",
              "<span style=\"color: #008000; font-weight: bold\">class</span> <span style=\"color: #00F; font-weight: bold\">Module</span>:\n",
              "    <span style=\"color: #A2F\">@T</span><span style=\"color: #A2F; font-weight: bold\">.</span>prim_func\n",
              "    <span style=\"color: #008000; font-weight: bold\">def</span> <span style=\"color: #00F\">add</span>(A: T<span style=\"color: #A2F; font-weight: bold\">.</span>Buffer((<span style=\"color: #008000\">16</span>, <span style=\"color: #008000\">64</span>, <span style=\"color: #008000\">1</span>, <span style=\"color: #008000\">16</span>), <span style=\"color: #BA2121\">&quot;int32&quot;</span>), B: T<span style=\"color: #A2F; font-weight: bold\">.</span>Buffer((<span style=\"color: #008000\">16</span>, <span style=\"color: #008000\">64</span>, <span style=\"color: #008000\">1</span>, <span style=\"color: #008000\">16</span>), <span style=\"color: #BA2121\">&quot;int32&quot;</span>), C: T<span style=\"color: #A2F; font-weight: bold\">.</span>Buffer((<span style=\"color: #008000\">16</span>, <span style=\"color: #008000\">64</span>, <span style=\"color: #008000\">1</span>, <span style=\"color: #008000\">16</span>), <span style=\"color: #BA2121\">&quot;int8&quot;</span>)):\n",
              "        T<span style=\"color: #A2F; font-weight: bold\">.</span>func_attr({<span style=\"color: #BA2121\">&quot;tir.noalias&quot;</span>: T<span style=\"color: #A2F; font-weight: bold\">.</span>bool(<span style=\"color: #008000; font-weight: bold\">True</span>)})\n",
              "        <span style=\"color: #007979; font-style: italic\"># with T.block(&quot;root&quot;):</span>\n",
              "        A_buf <span style=\"color: #A2F; font-weight: bold\">=</span> T<span style=\"color: #A2F; font-weight: bold\">.</span>alloc_buffer((<span style=\"color: #008000\">16</span>, <span style=\"color: #008000\">64</span>, <span style=\"color: #008000\">1</span>, <span style=\"color: #008000\">16</span>), <span style=\"color: #BA2121\">&quot;int32&quot;</span>)\n",
              "        B_buf <span style=\"color: #A2F; font-weight: bold\">=</span> T<span style=\"color: #A2F; font-weight: bold\">.</span>alloc_buffer((<span style=\"color: #008000\">16</span>, <span style=\"color: #008000\">64</span>, <span style=\"color: #008000\">1</span>, <span style=\"color: #008000\">16</span>), <span style=\"color: #BA2121\">&quot;int32&quot;</span>)\n",
              "        C_buf <span style=\"color: #A2F; font-weight: bold\">=</span> T<span style=\"color: #A2F; font-weight: bold\">.</span>alloc_buffer((<span style=\"color: #008000\">16</span>, <span style=\"color: #008000\">64</span>, <span style=\"color: #008000\">1</span>, <span style=\"color: #008000\">16</span>), <span style=\"color: #BA2121\">&quot;int32&quot;</span>)\n",
              "        <span style=\"color: #008000; font-weight: bold\">for</span> i0, i1, i2, i3 <span style=\"color: #008000; font-weight: bold\">in</span> T<span style=\"color: #A2F; font-weight: bold\">.</span>grid(<span style=\"color: #008000\">16</span>, <span style=\"color: #008000\">64</span>, <span style=\"color: #008000\">1</span>, <span style=\"color: #008000\">16</span>):\n",
              "            <span style=\"color: #008000; font-weight: bold\">with</span> T<span style=\"color: #A2F; font-weight: bold\">.</span>block(<span style=\"color: #BA2121\">&quot;A_buf&quot;</span>):\n",
              "                v_i0, v_i1, v_i2, v_i3 <span style=\"color: #A2F; font-weight: bold\">=</span> T<span style=\"color: #A2F; font-weight: bold\">.</span>axis<span style=\"color: #A2F; font-weight: bold\">.</span>remap(<span style=\"color: #BA2121\">&quot;SSSS&quot;</span>, [i0, i1, i2, i3])\n",
              "                T<span style=\"color: #A2F; font-weight: bold\">.</span>reads(A[v_i0, v_i1, v_i2, v_i3])\n",
              "                T<span style=\"color: #A2F; font-weight: bold\">.</span>writes(A_buf[v_i0, v_i1, v_i2, v_i3])\n",
              "                A_buf[v_i0, v_i1, v_i2, v_i3] <span style=\"color: #A2F; font-weight: bold\">=</span> A[v_i0, v_i1, v_i2, v_i3]\n",
              "        <span style=\"color: #008000; font-weight: bold\">for</span> i0, i1, i2, i3 <span style=\"color: #008000; font-weight: bold\">in</span> T<span style=\"color: #A2F; font-weight: bold\">.</span>grid(<span style=\"color: #008000\">16</span>, <span style=\"color: #008000\">64</span>, <span style=\"color: #008000\">1</span>, <span style=\"color: #008000\">16</span>):\n",
              "            <span style=\"color: #008000; font-weight: bold\">with</span> T<span style=\"color: #A2F; font-weight: bold\">.</span>block(<span style=\"color: #BA2121\">&quot;B_buf&quot;</span>):\n",
              "                v_i0, v_i1, v_i2, v_i3 <span style=\"color: #A2F; font-weight: bold\">=</span> T<span style=\"color: #A2F; font-weight: bold\">.</span>axis<span style=\"color: #A2F; font-weight: bold\">.</span>remap(<span style=\"color: #BA2121\">&quot;SSSS&quot;</span>, [i0, i1, i2, i3])\n",
              "                T<span style=\"color: #A2F; font-weight: bold\">.</span>reads(B[v_i0, v_i1, v_i2, v_i3])\n",
              "                T<span style=\"color: #A2F; font-weight: bold\">.</span>writes(B_buf[v_i0, v_i1, v_i2, v_i3])\n",
              "                B_buf[v_i0, v_i1, v_i2, v_i3] <span style=\"color: #A2F; font-weight: bold\">=</span> B[v_i0, v_i1, v_i2, v_i3]\n",
              "        <span style=\"color: #008000; font-weight: bold\">for</span> i0, i1, i2, i3 <span style=\"color: #008000; font-weight: bold\">in</span> T<span style=\"color: #A2F; font-weight: bold\">.</span>grid(<span style=\"color: #008000\">16</span>, <span style=\"color: #008000\">64</span>, <span style=\"color: #008000\">1</span>, <span style=\"color: #008000\">16</span>):\n",
              "            <span style=\"color: #008000; font-weight: bold\">with</span> T<span style=\"color: #A2F; font-weight: bold\">.</span>block(<span style=\"color: #BA2121\">&quot;C_buf&quot;</span>):\n",
              "                v_i0, v_i1, v_i2, v_i3 <span style=\"color: #A2F; font-weight: bold\">=</span> T<span style=\"color: #A2F; font-weight: bold\">.</span>axis<span style=\"color: #A2F; font-weight: bold\">.</span>remap(<span style=\"color: #BA2121\">&quot;SSSS&quot;</span>, [i0, i1, i2, i3])\n",
              "                T<span style=\"color: #A2F; font-weight: bold\">.</span>reads(A_buf[v_i0, v_i1, v_i2, v_i3], B_buf[v_i0, v_i1, v_i2, v_i3])\n",
              "                T<span style=\"color: #A2F; font-weight: bold\">.</span>writes(C_buf[v_i0, v_i1, v_i2, v_i3])\n",
              "                C_buf[v_i0, v_i1, v_i2, v_i3] <span style=\"color: #A2F; font-weight: bold\">=</span> A_buf[v_i0, v_i1, v_i2, v_i3] <span style=\"color: #A2F; font-weight: bold\">+</span> B_buf[v_i0, v_i1, v_i2, v_i3]\n",
              "        <span style=\"color: #008000; font-weight: bold\">for</span> i0, i1, i2, i3 <span style=\"color: #008000; font-weight: bold\">in</span> T<span style=\"color: #A2F; font-weight: bold\">.</span>grid(<span style=\"color: #008000\">16</span>, <span style=\"color: #008000\">64</span>, <span style=\"color: #008000\">1</span>, <span style=\"color: #008000\">16</span>):\n",
              "            <span style=\"color: #008000; font-weight: bold\">with</span> T<span style=\"color: #A2F; font-weight: bold\">.</span>block(<span style=\"color: #BA2121\">&quot;C&quot;</span>):\n",
              "                v_i0, v_i1, v_i2, v_i3 <span style=\"color: #A2F; font-weight: bold\">=</span> T<span style=\"color: #A2F; font-weight: bold\">.</span>axis<span style=\"color: #A2F; font-weight: bold\">.</span>remap(<span style=\"color: #BA2121\">&quot;SSSS&quot;</span>, [i0, i1, i2, i3])\n",
              "                T<span style=\"color: #A2F; font-weight: bold\">.</span>reads(C_buf[v_i0, v_i1, v_i2, v_i3])\n",
              "                T<span style=\"color: #A2F; font-weight: bold\">.</span>writes(C[v_i0, v_i1, v_i2, v_i3])\n",
              "                C[v_i0, v_i1, v_i2, v_i3] <span style=\"color: #A2F; font-weight: bold\">=</span> T<span style=\"color: #A2F; font-weight: bold\">.</span>Cast(<span style=\"color: #BA2121\">&quot;int8&quot;</span>, C_buf[v_i0, v_i1, v_i2, v_i3])\n",
              "</pre></div>\n"
            ],
            "text/plain": [
              "<IPython.core.display.HTML object>"
            ]
          },
          "metadata": {},
          "output_type": "display_data"
        }
      ],
      "source": [
        "func_name = \"add\"\n",
        "te_func = te.create_prim_func([A, B, C]).with_attr({\"global_symbol\": func_name})\n",
        "MyModule = IRModule({func_name: te_func})\n",
        "sch = tvm.tir.Schedule(MyModule)\n",
        "sch.mod.show()"
      ]
    },
    {
      "cell_type": "code",
      "execution_count": 17,
      "metadata": {},
      "outputs": [],
      "source": [
        "s = te.create_schedule(C.op)"
      ]
    },
    {
      "cell_type": "markdown",
      "metadata": {},
      "source": [
        "虽然此调度是合理的，但它不会编译到 VTA。为了获得正确的代码生成（code generation），需要应用调度原语（scheduling primitives）和代码注解（code annotation），将调度变换为可以直接 lower 到 VTA 硬件 intrinsics。其中包括：\n",
        "\n",
        "- DMA copy 运算将把全局作用域的张量复制到局部作用域的张量。\n",
        "- 执行向量加法的向量 ALU 运算。\n",
        "\n",
        "### Buffer 作用域\n",
        "\n",
        "首先，设置复制 buffer 的作用域，以指示 TVM 这些中间张量将存储在 VTA 的 on-chip SRAM buffer 中。下面，告诉 TVM `A_buf`、`B_buf`、`C_buf` 将存在于 VTA 的 on-chip *accumulator buffer* 中，该 buffer 作为 VTA 的通用寄存器（register）文件。\n",
        "\n",
        "将中间张量的作用域设置为 VTA 的 on-chip accumulator buffer"
      ]
    },
    {
      "cell_type": "code",
      "execution_count": 18,
      "metadata": {
        "collapsed": false
      },
      "outputs": [
        {
          "data": {
            "text/html": [
              "<div class=\"highlight\" style=\"background: \"><pre style=\"line-height: 125%;\"><span></span><span style=\"color: #007979; font-style: italic\"># from tvm.script import ir as I</span>\n",
              "<span style=\"color: #007979; font-style: italic\"># from tvm.script import tir as T</span>\n",
              "\n",
              "<span style=\"color: #A2F\">@I</span><span style=\"color: #A2F; font-weight: bold\">.</span>ir_module\n",
              "<span style=\"color: #008000; font-weight: bold\">class</span> <span style=\"color: #00F; font-weight: bold\">Module</span>:\n",
              "    <span style=\"color: #A2F\">@T</span><span style=\"color: #A2F; font-weight: bold\">.</span>prim_func\n",
              "    <span style=\"color: #008000; font-weight: bold\">def</span> <span style=\"color: #00F\">main</span>(A: T<span style=\"color: #A2F; font-weight: bold\">.</span>Buffer((<span style=\"color: #008000\">16</span>, <span style=\"color: #008000\">64</span>, <span style=\"color: #008000\">1</span>, <span style=\"color: #008000\">16</span>), <span style=\"color: #BA2121\">&quot;int32&quot;</span>), B: T<span style=\"color: #A2F; font-weight: bold\">.</span>Buffer((<span style=\"color: #008000\">16</span>, <span style=\"color: #008000\">64</span>, <span style=\"color: #008000\">1</span>, <span style=\"color: #008000\">16</span>), <span style=\"color: #BA2121\">&quot;int32&quot;</span>), C: T<span style=\"color: #A2F; font-weight: bold\">.</span>Buffer((<span style=\"color: #008000\">16</span>, <span style=\"color: #008000\">64</span>, <span style=\"color: #008000\">1</span>, <span style=\"color: #008000\">16</span>), <span style=\"color: #BA2121\">&quot;int8&quot;</span>)):\n",
              "        T<span style=\"color: #A2F; font-weight: bold\">.</span>func_attr({<span style=\"color: #BA2121\">&quot;from_legacy_te_schedule&quot;</span>: T<span style=\"color: #A2F; font-weight: bold\">.</span>bool(<span style=\"color: #008000; font-weight: bold\">True</span>), <span style=\"color: #BA2121\">&quot;tir.noalias&quot;</span>: T<span style=\"color: #A2F; font-weight: bold\">.</span>bool(<span style=\"color: #008000; font-weight: bold\">True</span>)})\n",
              "        A_buf <span style=\"color: #A2F; font-weight: bold\">=</span> T<span style=\"color: #A2F; font-weight: bold\">.</span>allocate([<span style=\"color: #008000\">32768</span>], <span style=\"color: #BA2121\">&quot;int32&quot;</span>, <span style=\"color: #BA2121\">&quot;local.acc_buffer&quot;</span>)\n",
              "        A_buf_1 <span style=\"color: #A2F; font-weight: bold\">=</span> T<span style=\"color: #A2F; font-weight: bold\">.</span>Buffer((<span style=\"color: #008000\">16384</span>,), <span style=\"color: #BA2121\">&quot;int32&quot;</span>, data<span style=\"color: #A2F; font-weight: bold\">=</span>A_buf, scope<span style=\"color: #A2F; font-weight: bold\">=</span><span style=\"color: #BA2121\">&quot;local.acc_buffer&quot;</span>, align<span style=\"color: #A2F; font-weight: bold\">=</span><span style=\"color: #008000\">16</span>)\n",
              "        <span style=\"color: #008000; font-weight: bold\">for</span> i0, i1, i3 <span style=\"color: #008000; font-weight: bold\">in</span> T<span style=\"color: #A2F; font-weight: bold\">.</span>grid(<span style=\"color: #008000\">16</span>, <span style=\"color: #008000\">64</span>, <span style=\"color: #008000\">16</span>):\n",
              "            cse_var_1: T<span style=\"color: #A2F; font-weight: bold\">.</span>int32 <span style=\"color: #A2F; font-weight: bold\">=</span> i0 <span style=\"color: #A2F; font-weight: bold\">*</span> <span style=\"color: #008000\">1024</span> <span style=\"color: #A2F; font-weight: bold\">+</span> i1 <span style=\"color: #A2F; font-weight: bold\">*</span> <span style=\"color: #008000\">16</span> <span style=\"color: #A2F; font-weight: bold\">+</span> i3\n",
              "            A_1 <span style=\"color: #A2F; font-weight: bold\">=</span> T<span style=\"color: #A2F; font-weight: bold\">.</span>Buffer((<span style=\"color: #008000\">16384</span>,), <span style=\"color: #BA2121\">&quot;int32&quot;</span>, data<span style=\"color: #A2F; font-weight: bold\">=</span>A<span style=\"color: #A2F; font-weight: bold\">.</span>data)\n",
              "            A_buf_1[cse_var_1] <span style=\"color: #A2F; font-weight: bold\">=</span> A_1[cse_var_1]\n",
              "        A_buf_2 <span style=\"color: #A2F; font-weight: bold\">=</span> T<span style=\"color: #A2F; font-weight: bold\">.</span>Buffer((<span style=\"color: #008000\">16384</span>,), <span style=\"color: #BA2121\">&quot;int32&quot;</span>, data<span style=\"color: #A2F; font-weight: bold\">=</span>A_buf, scope<span style=\"color: #A2F; font-weight: bold\">=</span><span style=\"color: #BA2121\">&quot;local.acc_buffer&quot;</span>, align<span style=\"color: #A2F; font-weight: bold\">=</span><span style=\"color: #008000\">16</span>)\n",
              "        <span style=\"color: #008000; font-weight: bold\">for</span> i0, i1, i3 <span style=\"color: #008000; font-weight: bold\">in</span> T<span style=\"color: #A2F; font-weight: bold\">.</span>grid(<span style=\"color: #008000\">16</span>, <span style=\"color: #008000\">64</span>, <span style=\"color: #008000\">16</span>):\n",
              "            cse_var_2: T<span style=\"color: #A2F; font-weight: bold\">.</span>int32 <span style=\"color: #A2F; font-weight: bold\">=</span> i0 <span style=\"color: #A2F; font-weight: bold\">*</span> <span style=\"color: #008000\">1024</span> <span style=\"color: #A2F; font-weight: bold\">+</span> i1 <span style=\"color: #A2F; font-weight: bold\">*</span> <span style=\"color: #008000\">16</span> <span style=\"color: #A2F; font-weight: bold\">+</span> i3\n",
              "            B_1 <span style=\"color: #A2F; font-weight: bold\">=</span> T<span style=\"color: #A2F; font-weight: bold\">.</span>Buffer((<span style=\"color: #008000\">16384</span>,), <span style=\"color: #BA2121\">&quot;int32&quot;</span>, data<span style=\"color: #A2F; font-weight: bold\">=</span>B<span style=\"color: #A2F; font-weight: bold\">.</span>data)\n",
              "            A_buf_2[cse_var_2 <span style=\"color: #A2F; font-weight: bold\">+</span> <span style=\"color: #008000\">16384</span>] <span style=\"color: #A2F; font-weight: bold\">=</span> B_1[cse_var_2]\n",
              "        A_buf_3 <span style=\"color: #A2F; font-weight: bold\">=</span> T<span style=\"color: #A2F; font-weight: bold\">.</span>Buffer((<span style=\"color: #008000\">16384</span>,), <span style=\"color: #BA2121\">&quot;int32&quot;</span>, data<span style=\"color: #A2F; font-weight: bold\">=</span>A_buf, scope<span style=\"color: #A2F; font-weight: bold\">=</span><span style=\"color: #BA2121\">&quot;local.acc_buffer&quot;</span>, align<span style=\"color: #A2F; font-weight: bold\">=</span><span style=\"color: #008000\">16</span>)\n",
              "        <span style=\"color: #008000; font-weight: bold\">for</span> i0, i1, i3 <span style=\"color: #008000; font-weight: bold\">in</span> T<span style=\"color: #A2F; font-weight: bold\">.</span>grid(<span style=\"color: #008000\">16</span>, <span style=\"color: #008000\">64</span>, <span style=\"color: #008000\">16</span>):\n",
              "            cse_var_3: T<span style=\"color: #A2F; font-weight: bold\">.</span>int32 <span style=\"color: #A2F; font-weight: bold\">=</span> i0 <span style=\"color: #A2F; font-weight: bold\">*</span> <span style=\"color: #008000\">1024</span> <span style=\"color: #A2F; font-weight: bold\">+</span> i1 <span style=\"color: #A2F; font-weight: bold\">*</span> <span style=\"color: #008000\">16</span> <span style=\"color: #A2F; font-weight: bold\">+</span> i3\n",
              "            A_buf_3[cse_var_3] <span style=\"color: #A2F; font-weight: bold\">=</span> A_buf_1[cse_var_3] <span style=\"color: #A2F; font-weight: bold\">+</span> A_buf_2[cse_var_3 <span style=\"color: #A2F; font-weight: bold\">+</span> <span style=\"color: #008000\">16384</span>]\n",
              "        <span style=\"color: #008000; font-weight: bold\">for</span> i0, i1, i3 <span style=\"color: #008000; font-weight: bold\">in</span> T<span style=\"color: #A2F; font-weight: bold\">.</span>grid(<span style=\"color: #008000\">16</span>, <span style=\"color: #008000\">64</span>, <span style=\"color: #008000\">16</span>):\n",
              "            cse_var_4: T<span style=\"color: #A2F; font-weight: bold\">.</span>int32 <span style=\"color: #A2F; font-weight: bold\">=</span> i0 <span style=\"color: #A2F; font-weight: bold\">*</span> <span style=\"color: #008000\">1024</span> <span style=\"color: #A2F; font-weight: bold\">+</span> i1 <span style=\"color: #A2F; font-weight: bold\">*</span> <span style=\"color: #008000\">16</span> <span style=\"color: #A2F; font-weight: bold\">+</span> i3\n",
              "            C_1 <span style=\"color: #A2F; font-weight: bold\">=</span> T<span style=\"color: #A2F; font-weight: bold\">.</span>Buffer((<span style=\"color: #008000\">16384</span>,), <span style=\"color: #BA2121\">&quot;int8&quot;</span>, data<span style=\"color: #A2F; font-weight: bold\">=</span>C<span style=\"color: #A2F; font-weight: bold\">.</span>data)\n",
              "            C_1[cse_var_4] <span style=\"color: #A2F; font-weight: bold\">=</span> T<span style=\"color: #A2F; font-weight: bold\">.</span>Cast(<span style=\"color: #BA2121\">&quot;int8&quot;</span>, A_buf_3[cse_var_4])\n",
              "</pre></div>\n"
            ],
            "text/plain": [
              "<IPython.core.display.HTML object>"
            ]
          },
          "metadata": {},
          "output_type": "display_data"
        }
      ],
      "source": [
        "s[A_buf].set_scope(env.acc_scope)\n",
        "s[B_buf].set_scope(env.acc_scope)\n",
        "s[C_buf].set_scope(env.acc_scope)\n",
        "tvm.lower(s, [A, B, C], simple_mode=True).show()"
      ]
    },
    {
      "cell_type": "markdown",
      "metadata": {},
      "source": [
        "### DMA 传输\n",
        "\n",
        "需要调度 DMA 传输，以便将存储在 DRAM 中的数据在 VTA 片上 buffer 之间来回移动。插入 `dma_copy` pragmas 来告诉编译器，复制运算将通过 DMA 批量执行，这在硬件加速器中很常见。\n",
        "\n",
        "使用 DMA pragma 标记 buffer 副本，将复制循环映射到 DMA transfer 运算："
      ]
    },
    {
      "cell_type": "code",
      "execution_count": 19,
      "metadata": {
        "collapsed": false
      },
      "outputs": [
        {
          "data": {
            "text/html": [
              "<div class=\"highlight\" style=\"background: \"><pre style=\"line-height: 125%;\"><span></span><span style=\"color: #007979; font-style: italic\"># from tvm.script import ir as I</span>\n",
              "<span style=\"color: #007979; font-style: italic\"># from tvm.script import tir as T</span>\n",
              "\n",
              "<span style=\"color: #A2F\">@I</span><span style=\"color: #A2F; font-weight: bold\">.</span>ir_module\n",
              "<span style=\"color: #008000; font-weight: bold\">class</span> <span style=\"color: #00F; font-weight: bold\">Module</span>:\n",
              "    <span style=\"color: #A2F\">@T</span><span style=\"color: #A2F; font-weight: bold\">.</span>prim_func\n",
              "    <span style=\"color: #008000; font-weight: bold\">def</span> <span style=\"color: #00F\">main</span>(A: T<span style=\"color: #A2F; font-weight: bold\">.</span>Buffer((<span style=\"color: #008000\">16</span>, <span style=\"color: #008000\">64</span>, <span style=\"color: #008000\">1</span>, <span style=\"color: #008000\">16</span>), <span style=\"color: #BA2121\">&quot;int32&quot;</span>), B: T<span style=\"color: #A2F; font-weight: bold\">.</span>Buffer((<span style=\"color: #008000\">16</span>, <span style=\"color: #008000\">64</span>, <span style=\"color: #008000\">1</span>, <span style=\"color: #008000\">16</span>), <span style=\"color: #BA2121\">&quot;int32&quot;</span>), C: T<span style=\"color: #A2F; font-weight: bold\">.</span>Buffer((<span style=\"color: #008000\">16</span>, <span style=\"color: #008000\">64</span>, <span style=\"color: #008000\">1</span>, <span style=\"color: #008000\">16</span>), <span style=\"color: #BA2121\">&quot;int8&quot;</span>)):\n",
              "        T<span style=\"color: #A2F; font-weight: bold\">.</span>func_attr({<span style=\"color: #BA2121\">&quot;from_legacy_te_schedule&quot;</span>: T<span style=\"color: #A2F; font-weight: bold\">.</span>bool(<span style=\"color: #008000; font-weight: bold\">True</span>), <span style=\"color: #BA2121\">&quot;tir.noalias&quot;</span>: T<span style=\"color: #A2F; font-weight: bold\">.</span>bool(<span style=\"color: #008000; font-weight: bold\">True</span>)})\n",
              "        A_buf <span style=\"color: #A2F; font-weight: bold\">=</span> T<span style=\"color: #A2F; font-weight: bold\">.</span>allocate([<span style=\"color: #008000\">32768</span>], <span style=\"color: #BA2121\">&quot;int32&quot;</span>, <span style=\"color: #BA2121\">&quot;local.acc_buffer&quot;</span>)\n",
              "        i0 <span style=\"color: #A2F; font-weight: bold\">=</span> T<span style=\"color: #A2F; font-weight: bold\">.</span>int32()\n",
              "        A_buf_1 <span style=\"color: #A2F; font-weight: bold\">=</span> T<span style=\"color: #A2F; font-weight: bold\">.</span>Buffer((<span style=\"color: #008000\">16384</span>,), <span style=\"color: #BA2121\">&quot;int32&quot;</span>, data<span style=\"color: #A2F; font-weight: bold\">=</span>A_buf, scope<span style=\"color: #A2F; font-weight: bold\">=</span><span style=\"color: #BA2121\">&quot;local.acc_buffer&quot;</span>, align<span style=\"color: #A2F; font-weight: bold\">=</span><span style=\"color: #008000\">16</span>)\n",
              "        <span style=\"color: #008000; font-weight: bold\">with</span> T<span style=\"color: #A2F; font-weight: bold\">.</span>attr(T<span style=\"color: #A2F; font-weight: bold\">.</span>iter_var(i0, <span style=\"color: #008000; font-weight: bold\">None</span>, <span style=\"color: #BA2121\">&quot;DataPar&quot;</span>, <span style=\"color: #BA2121\">&quot;&quot;</span>), <span style=\"color: #BA2121\">&quot;pragma_dma_copy&quot;</span>, <span style=\"color: #008000\">1</span>):\n",
              "            <span style=\"color: #008000; font-weight: bold\">for</span> i0, i1, i3 <span style=\"color: #008000; font-weight: bold\">in</span> T<span style=\"color: #A2F; font-weight: bold\">.</span>grid(<span style=\"color: #008000\">16</span>, <span style=\"color: #008000\">64</span>, <span style=\"color: #008000\">16</span>):\n",
              "                cse_var_1: T<span style=\"color: #A2F; font-weight: bold\">.</span>int32 <span style=\"color: #A2F; font-weight: bold\">=</span> i0 <span style=\"color: #A2F; font-weight: bold\">*</span> <span style=\"color: #008000\">1024</span> <span style=\"color: #A2F; font-weight: bold\">+</span> i1 <span style=\"color: #A2F; font-weight: bold\">*</span> <span style=\"color: #008000\">16</span> <span style=\"color: #A2F; font-weight: bold\">+</span> i3\n",
              "                A_1 <span style=\"color: #A2F; font-weight: bold\">=</span> T<span style=\"color: #A2F; font-weight: bold\">.</span>Buffer((<span style=\"color: #008000\">16384</span>,), <span style=\"color: #BA2121\">&quot;int32&quot;</span>, data<span style=\"color: #A2F; font-weight: bold\">=</span>A<span style=\"color: #A2F; font-weight: bold\">.</span>data)\n",
              "                A_buf_1[cse_var_1] <span style=\"color: #A2F; font-weight: bold\">=</span> A_1[cse_var_1]\n",
              "        i0_1 <span style=\"color: #A2F; font-weight: bold\">=</span> T<span style=\"color: #A2F; font-weight: bold\">.</span>int32()\n",
              "        A_buf_2 <span style=\"color: #A2F; font-weight: bold\">=</span> T<span style=\"color: #A2F; font-weight: bold\">.</span>Buffer((<span style=\"color: #008000\">16384</span>,), <span style=\"color: #BA2121\">&quot;int32&quot;</span>, data<span style=\"color: #A2F; font-weight: bold\">=</span>A_buf, scope<span style=\"color: #A2F; font-weight: bold\">=</span><span style=\"color: #BA2121\">&quot;local.acc_buffer&quot;</span>, align<span style=\"color: #A2F; font-weight: bold\">=</span><span style=\"color: #008000\">16</span>)\n",
              "        <span style=\"color: #008000; font-weight: bold\">with</span> T<span style=\"color: #A2F; font-weight: bold\">.</span>attr(T<span style=\"color: #A2F; font-weight: bold\">.</span>iter_var(i0_1, <span style=\"color: #008000; font-weight: bold\">None</span>, <span style=\"color: #BA2121\">&quot;DataPar&quot;</span>, <span style=\"color: #BA2121\">&quot;&quot;</span>), <span style=\"color: #BA2121\">&quot;pragma_dma_copy&quot;</span>, <span style=\"color: #008000\">1</span>):\n",
              "            <span style=\"color: #008000; font-weight: bold\">for</span> i0_1, i1, i3 <span style=\"color: #008000; font-weight: bold\">in</span> T<span style=\"color: #A2F; font-weight: bold\">.</span>grid(<span style=\"color: #008000\">16</span>, <span style=\"color: #008000\">64</span>, <span style=\"color: #008000\">16</span>):\n",
              "                cse_var_2: T<span style=\"color: #A2F; font-weight: bold\">.</span>int32 <span style=\"color: #A2F; font-weight: bold\">=</span> i0_1 <span style=\"color: #A2F; font-weight: bold\">*</span> <span style=\"color: #008000\">1024</span> <span style=\"color: #A2F; font-weight: bold\">+</span> i1 <span style=\"color: #A2F; font-weight: bold\">*</span> <span style=\"color: #008000\">16</span> <span style=\"color: #A2F; font-weight: bold\">+</span> i3\n",
              "                B_1 <span style=\"color: #A2F; font-weight: bold\">=</span> T<span style=\"color: #A2F; font-weight: bold\">.</span>Buffer((<span style=\"color: #008000\">16384</span>,), <span style=\"color: #BA2121\">&quot;int32&quot;</span>, data<span style=\"color: #A2F; font-weight: bold\">=</span>B<span style=\"color: #A2F; font-weight: bold\">.</span>data)\n",
              "                A_buf_2[cse_var_2 <span style=\"color: #A2F; font-weight: bold\">+</span> <span style=\"color: #008000\">16384</span>] <span style=\"color: #A2F; font-weight: bold\">=</span> B_1[cse_var_2]\n",
              "        A_buf_3 <span style=\"color: #A2F; font-weight: bold\">=</span> T<span style=\"color: #A2F; font-weight: bold\">.</span>Buffer((<span style=\"color: #008000\">16384</span>,), <span style=\"color: #BA2121\">&quot;int32&quot;</span>, data<span style=\"color: #A2F; font-weight: bold\">=</span>A_buf, scope<span style=\"color: #A2F; font-weight: bold\">=</span><span style=\"color: #BA2121\">&quot;local.acc_buffer&quot;</span>, align<span style=\"color: #A2F; font-weight: bold\">=</span><span style=\"color: #008000\">16</span>)\n",
              "        <span style=\"color: #008000; font-weight: bold\">for</span> i0_2, i1, i3 <span style=\"color: #008000; font-weight: bold\">in</span> T<span style=\"color: #A2F; font-weight: bold\">.</span>grid(<span style=\"color: #008000\">16</span>, <span style=\"color: #008000\">64</span>, <span style=\"color: #008000\">16</span>):\n",
              "            cse_var_3: T<span style=\"color: #A2F; font-weight: bold\">.</span>int32 <span style=\"color: #A2F; font-weight: bold\">=</span> i0_2 <span style=\"color: #A2F; font-weight: bold\">*</span> <span style=\"color: #008000\">1024</span> <span style=\"color: #A2F; font-weight: bold\">+</span> i1 <span style=\"color: #A2F; font-weight: bold\">*</span> <span style=\"color: #008000\">16</span> <span style=\"color: #A2F; font-weight: bold\">+</span> i3\n",
              "            A_buf_3[cse_var_3] <span style=\"color: #A2F; font-weight: bold\">=</span> A_buf_1[cse_var_3] <span style=\"color: #A2F; font-weight: bold\">+</span> A_buf_2[cse_var_3 <span style=\"color: #A2F; font-weight: bold\">+</span> <span style=\"color: #008000\">16384</span>]\n",
              "        i0_2 <span style=\"color: #A2F; font-weight: bold\">=</span> T<span style=\"color: #A2F; font-weight: bold\">.</span>int32()\n",
              "        T<span style=\"color: #A2F; font-weight: bold\">.</span>attr(T<span style=\"color: #A2F; font-weight: bold\">.</span>iter_var(i0_2, <span style=\"color: #008000; font-weight: bold\">None</span>, <span style=\"color: #BA2121\">&quot;DataPar&quot;</span>, <span style=\"color: #BA2121\">&quot;&quot;</span>), <span style=\"color: #BA2121\">&quot;pragma_dma_copy&quot;</span>, <span style=\"color: #008000\">1</span>)\n",
              "        <span style=\"color: #008000; font-weight: bold\">for</span> i0_2, i1, i3 <span style=\"color: #008000; font-weight: bold\">in</span> T<span style=\"color: #A2F; font-weight: bold\">.</span>grid(<span style=\"color: #008000\">16</span>, <span style=\"color: #008000\">64</span>, <span style=\"color: #008000\">16</span>):\n",
              "            cse_var_4: T<span style=\"color: #A2F; font-weight: bold\">.</span>int32 <span style=\"color: #A2F; font-weight: bold\">=</span> i0_2 <span style=\"color: #A2F; font-weight: bold\">*</span> <span style=\"color: #008000\">1024</span> <span style=\"color: #A2F; font-weight: bold\">+</span> i1 <span style=\"color: #A2F; font-weight: bold\">*</span> <span style=\"color: #008000\">16</span> <span style=\"color: #A2F; font-weight: bold\">+</span> i3\n",
              "            C_1 <span style=\"color: #A2F; font-weight: bold\">=</span> T<span style=\"color: #A2F; font-weight: bold\">.</span>Buffer((<span style=\"color: #008000\">16384</span>,), <span style=\"color: #BA2121\">&quot;int8&quot;</span>, data<span style=\"color: #A2F; font-weight: bold\">=</span>C<span style=\"color: #A2F; font-weight: bold\">.</span>data)\n",
              "            C_1[cse_var_4] <span style=\"color: #A2F; font-weight: bold\">=</span> T<span style=\"color: #A2F; font-weight: bold\">.</span>Cast(<span style=\"color: #BA2121\">&quot;int8&quot;</span>, A_buf_3[cse_var_4])\n",
              "</pre></div>\n"
            ],
            "text/plain": [
              "<IPython.core.display.HTML object>"
            ]
          },
          "metadata": {},
          "output_type": "display_data"
        }
      ],
      "source": [
        "s[A_buf].pragma(s[A_buf].op.axis[0], env.dma_copy)\n",
        "s[B_buf].pragma(s[B_buf].op.axis[0], env.dma_copy)\n",
        "s[C].pragma(s[C].op.axis[0], env.dma_copy)\n",
        "tvm.lower(s, [A, B, C], simple_mode=True).show()"
      ]
    },
    {
      "cell_type": "markdown",
      "metadata": {},
      "source": [
        "### ALU 运算\n",
        "\n",
        "VTA 有向量 ALU，可以在累加器 buffer 中对张量执行向量运算。为了告诉 TVM 给定的运算需要映射到 VTA 的 vector ALU，需要显式地用 `env.alu` pragma 标记 vector 加法循环。\n",
        "\n",
        "告诉 TVM 计算需要在 VTA 的向量 ALU 上执行："
      ]
    },
    {
      "cell_type": "code",
      "execution_count": 20,
      "metadata": {
        "collapsed": false
      },
      "outputs": [
        {
          "data": {
            "text/html": [
              "<div class=\"highlight\" style=\"background: \"><pre style=\"line-height: 125%;\"><span></span><span style=\"color: #007979; font-style: italic\"># from tvm.script import ir as I</span>\n",
              "<span style=\"color: #007979; font-style: italic\"># from tvm.script import tir as T</span>\n",
              "\n",
              "<span style=\"color: #A2F\">@I</span><span style=\"color: #A2F; font-weight: bold\">.</span>ir_module\n",
              "<span style=\"color: #008000; font-weight: bold\">class</span> <span style=\"color: #00F; font-weight: bold\">Module</span>:\n",
              "    <span style=\"color: #A2F\">@T</span><span style=\"color: #A2F; font-weight: bold\">.</span>prim_func\n",
              "    <span style=\"color: #008000; font-weight: bold\">def</span> <span style=\"color: #00F\">main</span>(A: T<span style=\"color: #A2F; font-weight: bold\">.</span>Buffer((<span style=\"color: #008000\">16</span>, <span style=\"color: #008000\">64</span>, <span style=\"color: #008000\">1</span>, <span style=\"color: #008000\">16</span>), <span style=\"color: #BA2121\">&quot;int32&quot;</span>), B: T<span style=\"color: #A2F; font-weight: bold\">.</span>Buffer((<span style=\"color: #008000\">16</span>, <span style=\"color: #008000\">64</span>, <span style=\"color: #008000\">1</span>, <span style=\"color: #008000\">16</span>), <span style=\"color: #BA2121\">&quot;int32&quot;</span>), C: T<span style=\"color: #A2F; font-weight: bold\">.</span>Buffer((<span style=\"color: #008000\">16</span>, <span style=\"color: #008000\">64</span>, <span style=\"color: #008000\">1</span>, <span style=\"color: #008000\">16</span>), <span style=\"color: #BA2121\">&quot;int8&quot;</span>)):\n",
              "        T<span style=\"color: #A2F; font-weight: bold\">.</span>func_attr({<span style=\"color: #BA2121\">&quot;from_legacy_te_schedule&quot;</span>: T<span style=\"color: #A2F; font-weight: bold\">.</span>bool(<span style=\"color: #008000; font-weight: bold\">True</span>), <span style=\"color: #BA2121\">&quot;tir.noalias&quot;</span>: T<span style=\"color: #A2F; font-weight: bold\">.</span>bool(<span style=\"color: #008000; font-weight: bold\">True</span>)})\n",
              "        A_buf <span style=\"color: #A2F; font-weight: bold\">=</span> T<span style=\"color: #A2F; font-weight: bold\">.</span>allocate([<span style=\"color: #008000\">32768</span>], <span style=\"color: #BA2121\">&quot;int32&quot;</span>, <span style=\"color: #BA2121\">&quot;local.acc_buffer&quot;</span>)\n",
              "        i0 <span style=\"color: #A2F; font-weight: bold\">=</span> T<span style=\"color: #A2F; font-weight: bold\">.</span>int32()\n",
              "        A_buf_1 <span style=\"color: #A2F; font-weight: bold\">=</span> T<span style=\"color: #A2F; font-weight: bold\">.</span>Buffer((<span style=\"color: #008000\">16384</span>,), <span style=\"color: #BA2121\">&quot;int32&quot;</span>, data<span style=\"color: #A2F; font-weight: bold\">=</span>A_buf, scope<span style=\"color: #A2F; font-weight: bold\">=</span><span style=\"color: #BA2121\">&quot;local.acc_buffer&quot;</span>, align<span style=\"color: #A2F; font-weight: bold\">=</span><span style=\"color: #008000\">16</span>)\n",
              "        <span style=\"color: #008000; font-weight: bold\">with</span> T<span style=\"color: #A2F; font-weight: bold\">.</span>attr(T<span style=\"color: #A2F; font-weight: bold\">.</span>iter_var(i0, <span style=\"color: #008000; font-weight: bold\">None</span>, <span style=\"color: #BA2121\">&quot;DataPar&quot;</span>, <span style=\"color: #BA2121\">&quot;&quot;</span>), <span style=\"color: #BA2121\">&quot;pragma_dma_copy&quot;</span>, <span style=\"color: #008000\">1</span>):\n",
              "            <span style=\"color: #008000; font-weight: bold\">for</span> i0, i1, i3 <span style=\"color: #008000; font-weight: bold\">in</span> T<span style=\"color: #A2F; font-weight: bold\">.</span>grid(<span style=\"color: #008000\">16</span>, <span style=\"color: #008000\">64</span>, <span style=\"color: #008000\">16</span>):\n",
              "                cse_var_1: T<span style=\"color: #A2F; font-weight: bold\">.</span>int32 <span style=\"color: #A2F; font-weight: bold\">=</span> i0 <span style=\"color: #A2F; font-weight: bold\">*</span> <span style=\"color: #008000\">1024</span> <span style=\"color: #A2F; font-weight: bold\">+</span> i1 <span style=\"color: #A2F; font-weight: bold\">*</span> <span style=\"color: #008000\">16</span> <span style=\"color: #A2F; font-weight: bold\">+</span> i3\n",
              "                A_1 <span style=\"color: #A2F; font-weight: bold\">=</span> T<span style=\"color: #A2F; font-weight: bold\">.</span>Buffer((<span style=\"color: #008000\">16384</span>,), <span style=\"color: #BA2121\">&quot;int32&quot;</span>, data<span style=\"color: #A2F; font-weight: bold\">=</span>A<span style=\"color: #A2F; font-weight: bold\">.</span>data)\n",
              "                A_buf_1[cse_var_1] <span style=\"color: #A2F; font-weight: bold\">=</span> A_1[cse_var_1]\n",
              "        i0_1 <span style=\"color: #A2F; font-weight: bold\">=</span> T<span style=\"color: #A2F; font-weight: bold\">.</span>int32()\n",
              "        A_buf_2 <span style=\"color: #A2F; font-weight: bold\">=</span> T<span style=\"color: #A2F; font-weight: bold\">.</span>Buffer((<span style=\"color: #008000\">16384</span>,), <span style=\"color: #BA2121\">&quot;int32&quot;</span>, data<span style=\"color: #A2F; font-weight: bold\">=</span>A_buf, scope<span style=\"color: #A2F; font-weight: bold\">=</span><span style=\"color: #BA2121\">&quot;local.acc_buffer&quot;</span>, align<span style=\"color: #A2F; font-weight: bold\">=</span><span style=\"color: #008000\">16</span>)\n",
              "        <span style=\"color: #008000; font-weight: bold\">with</span> T<span style=\"color: #A2F; font-weight: bold\">.</span>attr(T<span style=\"color: #A2F; font-weight: bold\">.</span>iter_var(i0_1, <span style=\"color: #008000; font-weight: bold\">None</span>, <span style=\"color: #BA2121\">&quot;DataPar&quot;</span>, <span style=\"color: #BA2121\">&quot;&quot;</span>), <span style=\"color: #BA2121\">&quot;pragma_dma_copy&quot;</span>, <span style=\"color: #008000\">1</span>):\n",
              "            <span style=\"color: #008000; font-weight: bold\">for</span> i0_1, i1, i3 <span style=\"color: #008000; font-weight: bold\">in</span> T<span style=\"color: #A2F; font-weight: bold\">.</span>grid(<span style=\"color: #008000\">16</span>, <span style=\"color: #008000\">64</span>, <span style=\"color: #008000\">16</span>):\n",
              "                cse_var_2: T<span style=\"color: #A2F; font-weight: bold\">.</span>int32 <span style=\"color: #A2F; font-weight: bold\">=</span> i0_1 <span style=\"color: #A2F; font-weight: bold\">*</span> <span style=\"color: #008000\">1024</span> <span style=\"color: #A2F; font-weight: bold\">+</span> i1 <span style=\"color: #A2F; font-weight: bold\">*</span> <span style=\"color: #008000\">16</span> <span style=\"color: #A2F; font-weight: bold\">+</span> i3\n",
              "                B_1 <span style=\"color: #A2F; font-weight: bold\">=</span> T<span style=\"color: #A2F; font-weight: bold\">.</span>Buffer((<span style=\"color: #008000\">16384</span>,), <span style=\"color: #BA2121\">&quot;int32&quot;</span>, data<span style=\"color: #A2F; font-weight: bold\">=</span>B<span style=\"color: #A2F; font-weight: bold\">.</span>data)\n",
              "                A_buf_2[cse_var_2 <span style=\"color: #A2F; font-weight: bold\">+</span> <span style=\"color: #008000\">16384</span>] <span style=\"color: #A2F; font-weight: bold\">=</span> B_1[cse_var_2]\n",
              "        i0_2 <span style=\"color: #A2F; font-weight: bold\">=</span> T<span style=\"color: #A2F; font-weight: bold\">.</span>int32()\n",
              "        A_buf_3 <span style=\"color: #A2F; font-weight: bold\">=</span> T<span style=\"color: #A2F; font-weight: bold\">.</span>Buffer((<span style=\"color: #008000\">16384</span>,), <span style=\"color: #BA2121\">&quot;int32&quot;</span>, data<span style=\"color: #A2F; font-weight: bold\">=</span>A_buf, scope<span style=\"color: #A2F; font-weight: bold\">=</span><span style=\"color: #BA2121\">&quot;local.acc_buffer&quot;</span>, align<span style=\"color: #A2F; font-weight: bold\">=</span><span style=\"color: #008000\">16</span>)\n",
              "        <span style=\"color: #008000; font-weight: bold\">with</span> T<span style=\"color: #A2F; font-weight: bold\">.</span>attr(T<span style=\"color: #A2F; font-weight: bold\">.</span>iter_var(i0_2, <span style=\"color: #008000; font-weight: bold\">None</span>, <span style=\"color: #BA2121\">&quot;DataPar&quot;</span>, <span style=\"color: #BA2121\">&quot;&quot;</span>), <span style=\"color: #BA2121\">&quot;pragma_alu&quot;</span>, <span style=\"color: #008000\">1</span>):\n",
              "            <span style=\"color: #008000; font-weight: bold\">for</span> i0_2, i1, i3 <span style=\"color: #008000; font-weight: bold\">in</span> T<span style=\"color: #A2F; font-weight: bold\">.</span>grid(<span style=\"color: #008000\">16</span>, <span style=\"color: #008000\">64</span>, <span style=\"color: #008000\">16</span>):\n",
              "                cse_var_3: T<span style=\"color: #A2F; font-weight: bold\">.</span>int32 <span style=\"color: #A2F; font-weight: bold\">=</span> i0_2 <span style=\"color: #A2F; font-weight: bold\">*</span> <span style=\"color: #008000\">1024</span> <span style=\"color: #A2F; font-weight: bold\">+</span> i1 <span style=\"color: #A2F; font-weight: bold\">*</span> <span style=\"color: #008000\">16</span> <span style=\"color: #A2F; font-weight: bold\">+</span> i3\n",
              "                A_buf_3[cse_var_3] <span style=\"color: #A2F; font-weight: bold\">=</span> A_buf_1[cse_var_3] <span style=\"color: #A2F; font-weight: bold\">+</span> A_buf_2[cse_var_3 <span style=\"color: #A2F; font-weight: bold\">+</span> <span style=\"color: #008000\">16384</span>]\n",
              "        i0_3 <span style=\"color: #A2F; font-weight: bold\">=</span> T<span style=\"color: #A2F; font-weight: bold\">.</span>int32()\n",
              "        T<span style=\"color: #A2F; font-weight: bold\">.</span>attr(T<span style=\"color: #A2F; font-weight: bold\">.</span>iter_var(i0_3, <span style=\"color: #008000; font-weight: bold\">None</span>, <span style=\"color: #BA2121\">&quot;DataPar&quot;</span>, <span style=\"color: #BA2121\">&quot;&quot;</span>), <span style=\"color: #BA2121\">&quot;pragma_dma_copy&quot;</span>, <span style=\"color: #008000\">1</span>)\n",
              "        <span style=\"color: #008000; font-weight: bold\">for</span> i0_3, i1, i3 <span style=\"color: #008000; font-weight: bold\">in</span> T<span style=\"color: #A2F; font-weight: bold\">.</span>grid(<span style=\"color: #008000\">16</span>, <span style=\"color: #008000\">64</span>, <span style=\"color: #008000\">16</span>):\n",
              "            cse_var_4: T<span style=\"color: #A2F; font-weight: bold\">.</span>int32 <span style=\"color: #A2F; font-weight: bold\">=</span> i0_3 <span style=\"color: #A2F; font-weight: bold\">*</span> <span style=\"color: #008000\">1024</span> <span style=\"color: #A2F; font-weight: bold\">+</span> i1 <span style=\"color: #A2F; font-weight: bold\">*</span> <span style=\"color: #008000\">16</span> <span style=\"color: #A2F; font-weight: bold\">+</span> i3\n",
              "            C_1 <span style=\"color: #A2F; font-weight: bold\">=</span> T<span style=\"color: #A2F; font-weight: bold\">.</span>Buffer((<span style=\"color: #008000\">16384</span>,), <span style=\"color: #BA2121\">&quot;int8&quot;</span>, data<span style=\"color: #A2F; font-weight: bold\">=</span>C<span style=\"color: #A2F; font-weight: bold\">.</span>data)\n",
              "            C_1[cse_var_4] <span style=\"color: #A2F; font-weight: bold\">=</span> T<span style=\"color: #A2F; font-weight: bold\">.</span>Cast(<span style=\"color: #BA2121\">&quot;int8&quot;</span>, A_buf_3[cse_var_4])\n",
              "</pre></div>\n"
            ],
            "text/plain": [
              "<IPython.core.display.HTML object>"
            ]
          },
          "metadata": {},
          "output_type": "display_data"
        }
      ],
      "source": [
        "s[C_buf].pragma(C_buf.op.axis[0], env.alu)\n",
        "# 查看最终的 schedule\n",
        "tvm.lower(s, [A, B, C], simple_mode=True).show()"
      ]
    },
    {
      "cell_type": "markdown",
      "metadata": {},
      "source": [
        "这就结束了本教程的调度部分。\n",
        "\n",
        "## TVM 计算\n",
        "\n",
        "在完成指定调度之后，可以将它编译成 TVM 函数。默认情况下，TVM 编译成可以直接从 python 调用的类型消除（type-erased）函数。\n",
        "\n",
        "在下面一行中，使用 {func}`tvm.build` 来创建函数。`build` 函数接受调度、函数的期望签名（包括输入和输出）以及想要编译的目标语言。"
      ]
    },
    {
      "cell_type": "code",
      "execution_count": 21,
      "metadata": {
        "collapsed": false
      },
      "outputs": [],
      "source": [
        "# ctx = tvm.target.Target(\"ext_dev\", host=env.target_host)\n",
        "target = \"ext_dev\"\n",
        "my_vadd = vta.build(s, [A, B, C], target=target, name=\"my_vadd\")"
      ]
    },
    {
      "cell_type": "markdown",
      "metadata": {},
      "source": [
        "### 保存 Module\n",
        "\n",
        "TVM 把模块保存到文件中，这样以后就可以加载回来了。这被称为提前编译（ahead-of-time compilation），可以节省一些编译时间。\n",
        "\n",
        "更重要的是，这允许在开发机器上交叉编译可执行文件，并通过 RPC 将其发送到 Pynq FPGA 板上执行。\n",
        "\n",
        "将编译后的模块写入 object 文件。"
      ]
    },
    {
      "cell_type": "code",
      "execution_count": 22,
      "metadata": {
        "collapsed": false
      },
      "outputs": [],
      "source": [
        "temp = tvm.contrib.utils.tempdir()\n",
        "my_vadd.save(temp.relpath(\"vadd.o\"))"
      ]
    },
    {
      "cell_type": "markdown",
      "metadata": {},
      "source": [
        "通过 RPC 发送可执行文件："
      ]
    },
    {
      "cell_type": "code",
      "execution_count": 23,
      "metadata": {},
      "outputs": [],
      "source": [
        "remote.upload(temp.relpath(\"vadd.o\"))"
      ]
    },
    {
      "cell_type": "markdown",
      "metadata": {},
      "source": [
        "### 载入 Module\n",
        "\n",
        "可以从文件系统加载编译后的模块来运行代码。"
      ]
    },
    {
      "cell_type": "code",
      "execution_count": 24,
      "metadata": {
        "collapsed": false
      },
      "outputs": [
        {
          "name": "stderr",
          "output_type": "stream",
          "text": [
            "2025-07-27 20:01:40.990 INFO load_module /tmp/tmph4ufaphg/vadd.o\n"
          ]
        }
      ],
      "source": [
        "f = remote.load_module(\"vadd.o\")"
      ]
    },
    {
      "attachments": {},
      "cell_type": "markdown",
      "metadata": {},
      "source": [
        "## 运行函数\n",
        "\n",
        "编译后的 TVM 函数使用简洁的 C API，可以被任何语言调用。\n",
        "\n",
        "TVM 用 python 提供了数组 API 来帮助快速测试和原型化。数组 API 是基于 [DLPack](https://github.com/dmlc/dlpack) 标准的。\n",
        "\n",
        "- 首先创建远程上下文（用于 Pynq 上的远程执行）。\n",
        "- 然后 `tvm.nd.array` 对数据进行相应的格式化。\n",
        "- `f()` 运行实际的计算。\n",
        "- `numpy()` 将结果数组以可解释的格式复制回来。\n",
        "\n",
        "获取远程设备的上下文："
      ]
    },
    {
      "cell_type": "code",
      "execution_count": 25,
      "metadata": {},
      "outputs": [],
      "source": [
        "ctx = remote.ext_dev(0)"
      ]
    },
    {
      "cell_type": "markdown",
      "metadata": {},
      "source": [
        "使用 {func}`tvm.nd.array` 将输入/输出数组格式化为 DLPack 标准："
      ]
    },
    {
      "cell_type": "code",
      "execution_count": 26,
      "metadata": {
        "collapsed": false
      },
      "outputs": [],
      "source": [
        "from tvm.topi.utils import get_const_tuple\n",
        "A_nd = tvm.nd.array(A_packed, ctx)\n",
        "B_nd = tvm.nd.array(B_packed, ctx)\n",
        "C_nd = tvm.nd.empty(get_const_tuple(C.shape), C.dtype, ctx)"
      ]
    },
    {
      "cell_type": "markdown",
      "metadata": {},
      "source": [
        "调用模块来执行计算："
      ]
    },
    {
      "cell_type": "code",
      "execution_count": 27,
      "metadata": {},
      "outputs": [],
      "source": [
        "f(A_nd, B_nd, C_nd)"
      ]
    },
    {
      "cell_type": "markdown",
      "metadata": {},
      "source": [
        "## 验证 Correctness\n",
        "\n",
        "使用 `numpy` 计算引用的结果，并断言矩阵乘法的输出确实是正确的："
      ]
    },
    {
      "cell_type": "code",
      "execution_count": 28,
      "metadata": {},
      "outputs": [
        {
          "data": {
            "text/plain": [
              "((16, 1024), (16, 1024))"
            ]
          },
          "execution_count": 28,
          "metadata": {},
          "output_type": "execute_result"
        }
      ],
      "source": [
        "A_orig.shape, B_orig.shape"
      ]
    },
    {
      "cell_type": "code",
      "execution_count": 29,
      "metadata": {
        "collapsed": false
      },
      "outputs": [
        {
          "name": "stdout",
          "output_type": "stream",
          "text": [
            "ALU 加法测试成功！\n"
          ]
        }
      ],
      "source": [
        "C_ref = (A_orig.astype(env.acc_dtype) + B_orig.astype(env.acc_dtype)).astype(C.dtype)\n",
        "C_ref = C_ref.reshape(m//env.BATCH, env.BATCH, n//env.BLOCK_OUT, env.BLOCK_OUT).transpose((0, 2, 1, 3))\n",
        "np.testing.assert_equal(C_ref, C_nd.numpy())\n",
        "print(\"ALU 加法测试成功！\")"
      ]
    },
    {
      "attachments": {},
      "cell_type": "markdown",
      "metadata": {},
      "source": [
        "## 获取 ALU 数据搬运情况"
      ]
    },
    {
      "cell_type": "code",
      "execution_count": 30,
      "metadata": {},
      "outputs": [
        {
          "name": "stderr",
          "output_type": "stream",
          "text": [
            "[20:01:41] /media/pc/data/board/arria10/lxw/tasks/tvm-new/src/runtime/profiling.cc:101: Warning: No timer implementation for ext_dev, using default timer instead. It may be inaccurate or have extra overhead.\n"
          ]
        }
      ],
      "source": [
        "time_f = f.time_evaluator(f.entry_name, ctx, number=20)\n",
        "if env.TARGET in [\"sim\", \"tsim\"]:\n",
        "    # Check if we're in local RPC mode (allows us to rebuild the\n",
        "    # runtime on the fly when varying the VTA designs)\n",
        "    local_rpc = int(os.environ.get(\"VTA_LOCAL_SIM_RPC\", \"0\"))\n",
        "    if local_rpc:\n",
        "        if env.TARGET == \"sim\":\n",
        "            remote.get_function(\"vta.simulator.profiler_clear\")()\n",
        "        else:\n",
        "            remote.get_function(\"vta.tsim.profiler_clear\")()\n",
        "        cost = time_f(A_nd, B_nd, C_nd)\n",
        "        if env.TARGET == \"sim\":\n",
        "            stats = json.loads(remote.get_function(\"vta.simulator.profiler_status\")())\n",
        "        else:\n",
        "            stats = json.loads(remote.get_function(\"vta.tsim.profiler_status\")())\n",
        "    else:\n",
        "        simulator.clear_stats()\n",
        "        cost = time_f(A_nd, B_nd, C_nd)\n",
        "        stats = simulator.stats()\n",
        "else:\n",
        "    cost = time_f(A_nd, B_nd, C_nd)"
      ]
    },
    {
      "cell_type": "code",
      "execution_count": 31,
      "metadata": {},
      "outputs": [
        {
          "data": {
            "text/plain": [
              "{'inp_load_nbytes': 0,\n",
              " 'wgt_load_nbytes': 0,\n",
              " 'acc_load_nbytes': 2752512,\n",
              " 'uop_load_nbytes': 84,\n",
              " 'out_store_nbytes': 344064,\n",
              " 'gemm_counter': 0,\n",
              " 'alu_counter': 21504}"
            ]
          },
          "execution_count": 31,
          "metadata": {},
          "output_type": "execute_result"
        }
      ],
      "source": [
        "stats"
      ]
    },
    {
      "cell_type": "code",
      "execution_count": 32,
      "metadata": {},
      "outputs": [
        {
          "data": {
            "text/plain": [
              "BenchmarkResult(min=8.322765e-05, mean=8.322765e-05, median=8.322765e-05, max=8.322765e-05, std=0.0, results=(8.322765e-05,))"
            ]
          },
          "execution_count": 32,
          "metadata": {},
          "output_type": "execute_result"
        }
      ],
      "source": [
        "cost"
      ]
    },
    {
      "cell_type": "markdown",
      "metadata": {},
      "source": [
        "```{topic} 小结\n",
        "\n",
        "本教程通过简单的向量加法示例，为深度学习加速器 VTA 编程提供了 TVM 演练。一般工作流程包括：\n",
        "\n",
        "- 用 VTA bitstream 在 RPC 上编程 FPGA。\n",
        "- 通过一系列的计算来描述向量加法的计算。\n",
        "- 描述如何使用调度原语执行计算。\n",
        "- 将函数编译到 VTA 目标。\n",
        "- 运行编译后的模块，并根据 `numpy` 实现来验证它。\n",
        "\n",
        "可以查看其他示例和教程，以了解更多有关 TVM 支持的运算、调度原语和其他功能，以编程 VTA。\n",
        "```"
      ]
    }
  ],
  "metadata": {
    "kernelspec": {
      "display_name": "Python 3",
      "language": "python",
      "name": "python3"
    },
    "language_info": {
      "codemirror_mode": {
        "name": "ipython",
        "version": 3
      },
      "file_extension": ".py",
      "mimetype": "text/x-python",
      "name": "python",
      "nbconvert_exporter": "python",
      "pygments_lexer": "ipython3",
      "version": "3.12.2"
    }
  },
  "nbformat": 4,
  "nbformat_minor": 0
}
